#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable
#define ACCUM_FLT4 float4
#define FLT float
#define FLT2 float2
#define FLT3 float3
#define FLT4 float4
#define TO_FLT4 convert_float4
#define TO_ACCUM_TYPE convert_float4
#define TO_ACCUM_FLT convert_float
#define READ_IMAGE read_imagef
#define WRITE_IMAGE write_imagef
__constant sampler_t smp_edge = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
__constant sampler_t smp_none = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;
__constant sampler_t smp_zero = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
// common definition
#define ACCUM_FLT float
constant ACCUM_FLT Bt[36] = {
        1.0000000000f, 0.0000000000f, -2.5000004768f, -0.0000001192f, 1.0000001192f, 0.0000000000f, 
        0.0000000000f, 0.9428091049f, 1.3333333731f, -0.4714044929f, -0.6666667461f, 0.0000000000f, 
        0.0000000000f, -0.9428089857f, 1.3333334923f, 0.4714045525f, -0.6666667461f, 0.0000000000f, 
        0.0000000000f, -0.1178511307f, -0.0833333358f, 0.2357022613f, 0.1666666865f, 0.0000000000f, 
        0.0000000000f, 0.1178511307f, -0.0833333507f, -0.2357022911f, 0.1666666865f, 0.0000000000f, 
        0.0000000000f, 0.9999998808f, -0.0000000596f, -2.5000000000f, 0.0000000000f, 1.0000000000f, 
};
// 6*6的bt，在texture2d中还有个6*8的
__kernel void winograd_4x4_to_36_read(
__global float4* dst_tensor_buffer,     // transformed input
  __global float4* src_tensor_buffer,   // src input
  __read_only image2d_t bt_tex2d,       // 6*8的那个bt
  int4 shared_int4_0,                   // 14400, 32, 120, -1 = 120*120, 128/4, 480/4, -padding
  int4 shared_int4_1,                   // 480, 480, -1, 36
  int4 shared_int4_2) {                 // 14400, 0, 0, 0
  int DST_X = get_global_id(0);         // 这个应该range(0, total_tiles)
  int DST_Y = get_global_id(1);         // 这个应该是6, bt矩阵的行数
  int DST_Z = get_global_id(2);         // 这个应该是range dst_channel_slice
                                        // 这个kernel每次是算一个slice中的一个tile(也即4个channel，h和w都为6的矩阵)和Bt中一列（1*6）的计算结果
                                        // 得到一个4*6*1(4个channel，6行1列的计算结果)

  //if (DST_X >= args.tiles_total || DST_Y >= 6 || DST_Z >= args.dst_tensor.Slices())
  if (DST_X >= shared_int4_0.x || DST_Y >= 6 || DST_Z >= shared_int4_0.y) { // 14400 32
    return; 
  }
  // 获取tile_x, tile_y的下标在input中的下标
  int tile_x = (DST_X % shared_int4_0.z) * 4;   // (DST_X % args.tiles_x) * 4 = (DST_X % 120) * 4
  int tile_y = (DST_X / shared_int4_0.z) * 4;   // (DST_X / args.tiles_x) * 4
  ACCUM_FLT4 I0, I1, I2, I3, I4, I5;
  ACCUM_FLT bt_ar[6];                           // DST_Y是range(0~6)，从bt_tex2d中读取两个float4
                                                // bt__tex2d是1*12的f4
  ACCUM_FLT4 t0 = TO_ACCUM_TYPE(read_imagef(bt_tex2d, smp_none, (int2)(DST_Y * 2 + 0, 0)));
  ACCUM_FLT4 t1 = TO_ACCUM_TYPE(read_imagef(bt_tex2d, smp_none, (int2)(DST_Y * 2 + 1, 0)));
  DST_Y *= 6;
  bt_ar[0] = t0.x;                              // 将读出来的float4存入这个bt_ar数组中
  bt_ar[1] = t0.y;
  bt_ar[2] = t0.z;
  bt_ar[3] = t0.w;
  bt_ar[4] = t1.x;
  bt_ar[5] = t1.y;

  // for循环6次
  int xc0 = tile_x + shared_int4_0.w + 0;       // tile_x - 1 + 0, 获取的是这个tile左上角的x坐标
  ACCUM_FLT m0_x = (ACCUM_FLT)(xc0 >= 0 && xc0 < shared_int4_1.x);  // xc0 >= 0 && xc0 < tensor_width, val为0 or 1
  bool inx0 = (xc0 >= 0 && xc0 < shared_int4_1.x);
  xc0 = clamp(xc0, 0, shared_int4_1.x - 1);     // clamp一下xc0到0~tensor_width
  // (dst_tensor_slice * 480 + 0） * 480 + xc0  // 这个xc0的float4的index
  int src_a_0 = (((DST_Z) * shared_int4_1.y + (0)) * shared_int4_1.x + (xc0));; 

  int xc1 = tile_x + shared_int4_0.w + 1;
  ACCUM_FLT m1_x = (ACCUM_FLT)(xc1 >= 0 && xc1 < shared_int4_1.x);
  bool inx1 = (xc1 >= 0 && xc1 < shared_int4_1.x);
  xc1 = clamp(xc1, 0, shared_int4_1.x - 1);
  int src_a_1 = (((DST_Z) * shared_int4_1.y + (0)) * shared_int4_1.x + (xc1));;

  int xc2 = tile_x + shared_int4_0.w + 2;
  ACCUM_FLT m2_x = (ACCUM_FLT)(xc2 >= 0 && xc2 < shared_int4_1.x);
  bool inx2 = (xc2 >= 0 && xc2 < shared_int4_1.x);
  xc2 = clamp(xc2, 0, shared_int4_1.x - 1);
  int src_a_2 = (((DST_Z) * shared_int4_1.y + (0)) * shared_int4_1.x + (xc2));;

  int xc3 = tile_x + shared_int4_0.w + 3;
  ACCUM_FLT m3_x = (ACCUM_FLT)(xc3 >= 0 && xc3 < shared_int4_1.x);
  bool inx3 = (xc3 >= 0 && xc3 < shared_int4_1.x);
  xc3 = clamp(xc3, 0, shared_int4_1.x - 1);
  int src_a_3 = (((DST_Z) * shared_int4_1.y + (0)) * shared_int4_1.x + (xc3));;

  int xc4 = tile_x + shared_int4_0.w + 4;
  ACCUM_FLT m4_x = (ACCUM_FLT)(xc4 >= 0 && xc4 < shared_int4_1.x);
  bool inx4 = (xc4 >= 0 && xc4 < shared_int4_1.x);
  xc4 = clamp(xc4, 0, shared_int4_1.x - 1);
  int src_a_4 = (((DST_Z) * shared_int4_1.y + (0)) * shared_int4_1.x + (xc4));;

  int xc5 = tile_x + shared_int4_0.w + 5;
  ACCUM_FLT m5_x = (ACCUM_FLT)(xc5 >= 0 && xc5 < shared_int4_1.x);
  bool inx5 = (xc5 >= 0 && xc5 < shared_int4_1.x);
  xc5 = clamp(xc5, 0, shared_int4_1.x - 1);
  int src_a_5 = (((DST_Z) * shared_int4_1.y + (0)) * shared_int4_1.x + (xc5));;

  {
    int yc = tile_y + shared_int4_1.z;  // tile的起始y坐标, tile_y - 1
    bool iny = (yc >= 0 && yc < shared_int4_1.y);   // 在有效的坐标内
    int offset = select(0, yc * shared_int4_1.x, iny);    // 如果在有效的坐标内，则offset = yc * 480, 否则offset = 0
    ACCUM_FLT bt = bt_ar[0] * (ACCUM_FLT)(iny);           // 如果在有效坐标内，bt = br_ar[0], 否则bt = 0
    ACCUM_FLT4 src0 = src_tensor_buffer[src_a_0 + offset] * m0_x; // 取出一个f4, 若x方向不在有效坐标内，则m0_x为0, 否则为1
    I0 = bt * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[src_a_1 + offset] * m1_x;
    I1 = bt * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[src_a_2 + offset] * m2_x;
    I2 = bt * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[src_a_3 + offset] * m3_x;
    I3 = bt * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[src_a_4 + offset] * m4_x;
    I4 = bt * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[src_a_5 + offset] * m5_x;
    I5 = bt * src5;
  }
  {
    int yc = tile_y + shared_int4_1.z + (1);
    bool iny = (yc >= 0 && yc < shared_int4_1.y);
    int offset = select(0, yc * shared_int4_1.x, iny);
    ACCUM_FLT bt = bt_ar[1] * (ACCUM_FLT)(iny);
    ACCUM_FLT4 src0 = src_tensor_buffer[src_a_0 + offset] * m0_x;
    I0 += bt * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[src_a_1 + offset] * m1_x;
    I1 += bt * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[src_a_2 + offset] * m2_x;
    I2 += bt * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[src_a_3 + offset] * m3_x;
    I3 += bt * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[src_a_4 + offset] * m4_x;
    I4 += bt * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[src_a_5 + offset] * m5_x;
    I5 += bt * src5;
  }
  {
    int yc = tile_y + shared_int4_1.z + (2);
    bool iny = (yc >= 0 && yc < shared_int4_1.y);
    int offset = select(0, yc * shared_int4_1.x, iny);
    ACCUM_FLT bt = bt_ar[2] * (ACCUM_FLT)(iny);
    ACCUM_FLT4 src0 = src_tensor_buffer[src_a_0 + offset] * m0_x;
    I0 += bt * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[src_a_1 + offset] * m1_x;
    I1 += bt * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[src_a_2 + offset] * m2_x;
    I2 += bt * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[src_a_3 + offset] * m3_x;
    I3 += bt * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[src_a_4 + offset] * m4_x;
    I4 += bt * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[src_a_5 + offset] * m5_x;
    I5 += bt * src5;
  }
  {
    int yc = tile_y + shared_int4_1.z + (3);
    bool iny = (yc >= 0 && yc < shared_int4_1.y);
    int offset = select(0, yc * shared_int4_1.x, iny);
    ACCUM_FLT bt = bt_ar[3] * (ACCUM_FLT)(iny);
    ACCUM_FLT4 src0 = src_tensor_buffer[src_a_0 + offset] * m0_x;
    I0 += bt * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[src_a_1 + offset] * m1_x;
    I1 += bt * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[src_a_2 + offset] * m2_x;
    I2 += bt * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[src_a_3 + offset] * m3_x;
    I3 += bt * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[src_a_4 + offset] * m4_x;
    I4 += bt * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[src_a_5 + offset] * m5_x;
    I5 += bt * src5;
  }
  {
    int yc = tile_y + shared_int4_1.z + (4);
    bool iny = (yc >= 0 && yc < shared_int4_1.y);
    int offset = select(0, yc * shared_int4_1.x, iny);
    ACCUM_FLT bt = bt_ar[4] * (ACCUM_FLT)(iny);
    ACCUM_FLT4 src0 = src_tensor_buffer[src_a_0 + offset] * m0_x;
    I0 += bt * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[src_a_1 + offset] * m1_x;
    I1 += bt * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[src_a_2 + offset] * m2_x;
    I2 += bt * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[src_a_3 + offset] * m3_x;
    I3 += bt * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[src_a_4 + offset] * m4_x;
    I4 += bt * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[src_a_5 + offset] * m5_x;
    I5 += bt * src5;
  }
  {
    int yc = tile_y + shared_int4_1.z + (5);
    bool iny = (yc >= 0 && yc < shared_int4_1.y);
    int offset = select(0, yc * shared_int4_1.x, iny);
    ACCUM_FLT bt = bt_ar[5] * (ACCUM_FLT)(iny);
    ACCUM_FLT4 src0 = src_tensor_buffer[src_a_0 + offset] * m0_x;
    I0 += bt * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[src_a_1 + offset] * m1_x;
    I1 += bt * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[src_a_2 + offset] * m2_x;
    I2 += bt * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[src_a_3 + offset] * m3_x;
    I3 += bt * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[src_a_4 + offset] * m4_x;
    I4 += bt * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[src_a_5 + offset] * m5_x;
    I5 += bt * src5;
  }
  // 到此算出了tile小矩阵和bt中一列的结果


  {
    // r0 = I0 * bt[0] + I1 * bt[1] + I2 * bt[2] + ... + I5 * bt[5]
    // 但是bt[0]为1, bt[1], bt[3], bt[5]均为0, 所以此处这几个省去了
    FLT4 r0 = TO_FLT4(I0 + Bt[2] * I2 + Bt[4] * I4);
    dst_tensor_buffer[(((DST_Z) * shared_int4_1.w + (DST_Y)) * shared_int4_2.x + (DST_X))] = r0;
;
    DST_Y++;
  }
  {
    FLT4 r0 = TO_FLT4(Bt[7] * I1 + Bt[8] * I2 + Bt[9] * I3 + Bt[10] * I4);
    dst_tensor_buffer[(((DST_Z) * shared_int4_1.w + (DST_Y)) * shared_int4_2.x + (DST_X))] = r0;
;
    DST_Y++;
  }
  {
    FLT4 r0 = TO_FLT4(Bt[13] * I1 + Bt[14] * I2 + Bt[15] * I3 + Bt[16] * I4);
    dst_tensor_buffer[(((DST_Z) * shared_int4_1.w + (DST_Y)) * shared_int4_2.x + (DST_X))] = r0;
;
    DST_Y++;
  }
  {
    FLT4 r0 = TO_FLT4(Bt[19] * I1 + Bt[20] * I2 + Bt[21] * I3 + Bt[22] * I4);
    dst_tensor_buffer[(((DST_Z) * shared_int4_1.w + (DST_Y)) * shared_int4_2.x + (DST_X))] = r0;
;
    DST_Y++;
  }
  {
    FLT4 r0 = TO_FLT4(Bt[25] * I1 + Bt[26] * I2 + Bt[27] * I3 + Bt[28] * I4);
    dst_tensor_buffer[(((DST_Z) * shared_int4_1.w + (DST_Y)) * shared_int4_2.x + (DST_X))] = r0;
;
    DST_Y++;
  }
  {
    FLT4 r0 = TO_FLT4(Bt[31] * I1 + Bt[33] * I3 + I5);
    dst_tensor_buffer[(((DST_Z) * shared_int4_1.w + (DST_Y)) * shared_int4_2.x + (DST_X))] = r0;
;
    DST_Y++;
  }
}